[DSA] Make the fused top-k exact on an overflowing threshold bin by xiaofei-zheng · Pull Request #37941 · sgl-project/sglang

xiaofei-zheng · 2026-09-04T06:24:22Z

Motivation

The coarse histogram bin comes from the top bits of the fp16 cast of each score, so one bin spans a quarter-binade. Phase 3 stages threshold-bin candidates under if (count_eq < kMaxNumTie) and phase 4 clips tie_count = min(equal_count, kMaxNumTie), so once a threshold bin holds more than 2048 elements handle_tie ranks whichever arrived first and the kernel returns a wrong selected-value multiset with no diagnostic. This is the shared DSA indexer top-k, serving every model routed through srt/layers/attention/dsa.

Found while profiling GLM-5.2 decode on MI355X: device-side counters over an agentic serving run enter the truncation path on ~0.04% of rows across all 8 ranks (556–598 of 1.1–1.5M per rank), largest threshold bin 4139 against the 2048 cap, none of those rows bit-identical. #35257 reports the same defect independently from the CUDA side — referenced, not closed.

Modifications

Refine the overflowing bin on the exact key instead of truncating it: up to four 8-bit radix passes over the order-preserving key from extract_exact_bin. It exits once the refined set fits the staging buffer, or once all 32 bits are consumed — the key is injective on fp32 bit patterns, so the survivors are bit-identical by then and truncating them is legitimate. A min/max reduction skips the refinement when the set is already bit-identical.

The fast path branches on equal_count > kMaxNumTie and is otherwise unchanged; kMaxNumTie, kBlockSize and the shared-memory footprint are untouched.

Folded in from #37942 (closed): the ROCm __launch_bounds__ occupancy argument now states the physical wave floor, since HIP reads it as waves per SIMD where CUDA reads it as blocks per SM. Codegen-neutral alone; it matters because the plausible wrong reading caps the allocator at 64 VGPRs and would spill at the 77 this change needs.

Accuracy Tests

Pre-existing cases draw from torch.randn and tolerate MAX_PERMIT_ERROR = 5, so they never reach the overflow and pass on the truncating kernel. This adds narrow, narrow_bin and tiny distributions that collapse a row into one coarse bin, on a register-path (6×8192) and a streaming-path (4×32768) shape, asserting exactly. Against the unfixed kernel:

suite	result on the unfixed kernel
278 pre-existing cases	278 passed
6 new narrow / narrow_bin / tiny	6 failed, 227–9198 wrong selections each
2 new `all_equal` controls	2 passed

all_equal overflows too, but bit-identically, so any subset is correct and it must keep passing — that asymmetry pins the fault to truncation rather than mis-binning. All 286 pass with this change; a separate 50-case gate moves 29/50 to 50/50.

Benchmarking and Profiling

MI355X (gfx950), weighted J over the captured GLM-5.2 decode shapes, graph-replayed:

	c4	c10
before	22.398 µs	29.393 µs
after	22.336 µs	29.367 µs

A paired decode A/B measures −0.50% / −0.27% step time against a 0.04–0.08% within-arm spread, i.e. parity. This is a correctness fix, not a speedup. Register usage goes 57 → 78 VGPRs with zero spills, and the LDS footprint is unchanged at 27392 B.

Per review feedback the refinement was cut from up to ten input passes to at most four (one to seed the first histogram, then one per round, each fusing the previous round's emit). Kernel time on MI355X, 200 iterations after warmup:

distribution	shape	before	after
`randn` (no overflow, hot path)	6×8192	5.604 µs	5.596 µs	−0.1%
`randn` (no overflow, hot path)	4×32768	12.507 µs	12.475 µs	−0.3%
`narrow_bin` (every row refines)	6×8192	34.825 µs	22.482 µs	−35.4%
`narrow_bin` (every row refines)	4×32768	75.559 µs	53.223 µs	−29.6%
`two_values` (worst case, 32 bits)	6×8192	53.509 µs	36.354 µs	−32.1%
`two_values` (worst case, 32 bits)	4×32768	168.130 µs	119.571 µs	−28.9%

The hot path is untouched; only rows that actually overflow reach the refinement.

Known gap

The CUDA-only TopKCluster path still truncates; refining there needs cluster-wide rather than block-local histogram and emit counters, so it is left for a follow-up. #35257 quotes only the register and streaming sites, so that third one is unreported. Related but distinct, #36807 covers the same class of truncation in the AOT topk.cu.

Checklist

Format your code according to the Code Formatting with Pre-Commit.
Add unit tests as outlined in the Running Unit Tests.
Update documentation / docstrings / example tutorials as needed, according to Writing Documentation.

Known gap

The CUDA-only TopKCluster path still truncates. Refining there is a different routine: the candidate set is split across kClusterSize ranks, so each radix pass needs a cluster-wide histogram all-reduce and a cluster-wide emit counter rather than the block-local ones used here. Left for a follow-up rather than grown into this PR — happy to open a tracking issue.

CI States

Latest PR Test (Base): ⏳ Run #35081117275
Latest PR Test (Extra): ❌ Run #35081117023
Latest PR Test (AMD ROCm 10): ❌ Run #35081117212

xiaofei-zheng · 2026-09-09T05:28:02Z

@HaiShaw could you add the run-ci label here when you get a chance?

The gate currently exits with Missing required label 'run-ci', so every test job is skipped and this PR has never actually been exercised by CI — only lint and the other label-independent checks have run.

For context: this is a correctness fix in the fused DSA indexer top-k, found while profiling our own GLM-5.2 decode runs on MI355X. When a threshold coarse bin holds more than kMaxNumTie (2048) candidates, the collect pass keeps whichever arrived first and the kernel silently returns a wrong top-k set; device-side counters put that on ~0.04% of rows on a real serving run, with the largest observed bin at 4139. #35257 and #36807 report related problems in this area. #37942 has been folded into this PR and closed, so this is now the only one to look at.

HaiShaw · 2026-09-09T09:45:23Z

@BBuf @yuan-luo @DarkSharpness please review

HaiShaw · 2026-09-09T09:46:33Z

/tag-and-rerun-ci

DarkSharpness

Some comments on performance

DarkSharpness · 2026-09-09T11:38:05Z

+    // Early-out for a bit-identical candidate set: the radix passes cannot
+    // separate it and would only re-derive what phase 3 already staged. One
+    // distinct exact key means the candidates are interchangeable, so any
+    // kMaxNumTie-subset is correct, including the one already in tie.values.
+    // Comparing min against max of the key is exact for any value, and costs
+    // the one scan that non-degenerate overflow rows pay on top.
+    {
+      uint32_t key_min = 0xFFFFFFFFu;
+      uint32_t key_max = 0u;
+      for_each_input(problem.in, problem.seq_len, [&](float val, uint32_t) {
+        if (val >= v_lo && val < v_hi) {
+          const auto key = extract_exact_bin(val);
+          key_min = min(key_min, key);
+          key_max = max(key_max, key);
+        }
+      });
+      // Reduce per-thread extrema: two atomics per thread, not per candidate.
+      if (tx == 0) {
+        handle->histogram[0][0] = 0xFFFFFFFFu;
+        handle->histogram[0][1] = 0u;
+      }
+      __syncthreads();
+      atomicMin(&handle->histogram[0][0], key_min);
+      atomicMax(&handle->histogram[0][1], key_max);
+      __syncthreads();
+      const bool bit_identical = handle->histogram[0][0] == handle->histogram[0][1];
+      __syncthreads();  // all threads read the scratch before the loop clears it
+      if (bit_identical) {
+        // equal_count > kMaxNumTie on entry, so phase 3 filled the whole buffer.
+        const auto above_count = smem->count_gt;
+        const auto remain_topk = above_count < topk ? topk - above_count : 0;
+        handle_tie(smem->tie.values, problem, above_count, kMaxNumTie, remain_topk, handle);
+        return;
+      }
+    }


Why do we need this pass? It's not common that all keys are equal in fp32. We should optimize for hot path. This actually adds to additional linear scan cost in normal pass.

DarkSharpness · 2026-09-09T11:39:27Z

+    for (uint32_t round = 0; round < 4 && cand_count > kMaxNumTie && remain > 0; ++round) {
+      const uint32_t shift = 24 - round * 8;
+
+      if (tx < kRadixSize) handle->histogram[0][tx] = 0;
+      __syncthreads();
+      for_each_input(problem.in, problem.seq_len, [&](float val, uint32_t) {
+        if (val >= v_lo && val < v_hi) {
+          const auto key = extract_exact_bin(val);
+          if ((key & mask) == prefix) atomicAdd(&handle->histogram[0][(key >> shift) & 0xFFu], 1);
+        }
+      });
+      __syncthreads();
+
+      refine_find_threshold(cand_count, remain, handle);
+      const auto match = handle->match;
+
+      for_each_input(problem.in, problem.seq_len, [&](float val, uint32_t idx) {
+        if (val >= v_lo && val < v_hi) {
+          const auto key = extract_exact_bin(val);
+          if ((key & mask) == prefix && ((key >> shift) & 0xFFu) > match.bin) {
+            const auto pos = atomicAdd(&smem->count_gt, 1);
+            if (pos < topk) [[likely]]
+              problem.emit(pos, idx);
+          }
+        }
+      });
+
+      prefix |= match.bin << shift;
+      mask |= 0xFFu << shift;
+      remain -= match.above_count;
+      cand_count = match.equal_count;
+      __syncthreads();  // `match` is read above and overwritten by the next pass
+    }


Is this efficient enough? This would involve 8 linear pass. Please reduce that to at most 4 passes.

DarkSharpness · 2026-09-09T11:40:32Z

+    if (tx == 0) smem->count_eq = 0;
+    __syncthreads();
+    for_each_input(problem.in, problem.seq_len, [&](float val, uint32_t idx) {
+      if (val >= v_lo && val < v_hi) {
+        const auto key = extract_exact_bin(val);
+        if ((key & mask) == prefix) {
+          const auto slot = atomicAdd(&smem->count_eq, 1);
+          if (slot < kMaxNumTie) smem->tie.values[slot] = {val, idx};
+        }
+      }
+    });


same, this passes should be eliminated i guess

DarkSharpness

Some comments on performance

xiaofei-zheng · 2026-09-10T07:28:53Z

@DarkSharpness thanks — all three addressed in c05fe16, and measured on MI355X.

The min/max early-out is gone. You are right that it was paying a scan on every refinement to serve a degenerate case. Without it the radix rounds still terminate correctly on an all-equal row: they consume all 32 bits and then truncate a bit-identical set, which is exact. Removing it also freed the tie staging buffer, which this path does not need — it re-derives the candidates rather than trusting what the collect pass staged there — so the buffer now carries the histogram.

Down to at most four passes. Two changes got there. The radix is 12 bits instead of 8, using that freed buffer (it is exactly kRefineSize uint32 wide, so the shared-memory footprint is unchanged at 27392 B), which covers the 32-bit key in three rounds instead of four. And each round's emit is fused with the next round's histogram: one scan classifies a candidate as above / equal / below the round's threshold digit and does both. So the cost is one pass to seed the first histogram plus one per round — four worst case, two if the bin resolves in the first round.

The staging pass is eliminated too, fused into the last round's emit the same way.

Kernel time, 200 iterations after warmup:

distribution	shape	before	after
`randn` (no overflow, hot path)	6×8192	5.604 µs	5.596 µs	−0.1%
`randn` (no overflow, hot path)	4×32768	12.507 µs	12.475 µs	−0.3%
`narrow_bin` (every row refines)	6×8192	34.825 µs	22.482 µs	−35.4%
`narrow_bin` (every row refines)	4×32768	75.559 µs	53.223 µs	−29.6%
`two_values` (worst case, 32 bits)	6×8192	53.509 µs	36.354 µs	−32.1%
`two_values` (worst case, 32 bits)	4×32768	168.130 µs	119.571 µs	−28.9%

The hot path is unchanged, as it should be — it still only pays the equal_count > kMaxNumTie comparison.

On correctness after the rewrite: the suite is 288/288 on this branch. Run against the unfixed kernel on main it is 8 failed / 2 passed, the two passes being the all_equal controls, so the tests still discriminate. two_values is new and covers the exit the early-out used to stand in front of: two adjacent fp32 values at roughly half the row each, so the refinement has to consume all 32 bits and still has more than kMaxNumTie survivors at that point. Register usage is 78 VGPRs with zero spills, up from 57 on main.

xiaofei-zheng · 2026-09-11T02:44:15Z

@DarkSharpness when you have a moment — the performance changes you asked for are in fcd39053bd, and CI is now green on the jobs that cover this kernel.

The branch has also been rebased onto latest main. That turned out to matter: the CI maintenance-mode gate was rejecting the old base (Required base commit 3700c4ee26a1 ... your PR is diverged), which failed check-changes on the PR Test and PR Test Extra workflows and silently skipped the whole NVIDIA pipeline. So until the rebase this change had only ever been exercised on AMD. Both sides now run:


`call-jit-kernel-tests / jit-kernel-unit-test (0)`	pass — includes `test_topk_v2.py`
`call-jit-kernel-tests / jit-kernel-unit-test (1)`	pass
`call-jit-kernel-tests / jit-kernel-b200-test`	pass
`call-jit-kernel-tests / jit-kernel-multigpu-unit-test`	pass
`call-jit-kernel-tests / jit-kernel-benchmark-test`	pass
`jit-kernel-unit-test-amd (rocm10)`	pass — includes `test_topk_v2.py`
`jit-kernel-benchmark-test-amd (rocm10)`	pass

The remaining red checks are unrelated to this diff, as far as I can tell:

base-b-test-1-gpu-small (1..3) time out after 30 min on test_spec_eagle_topk.py (TestEagle3Topk16.test_acc_length), which brings up a Llama-3.1-8B + EAGLE3 server and carries est_time=876. The topk there is speculative_eagle_topk, not the DSA indexer.
stage-a-test-1-gpu-small-amd fails in test_wave_attention_kernels.py: grouped decode attention returns nan, the cos_sim > 0.99 assertion trips and the process then aborts on glibc heap corruption. That job runs three test files and none of them is a top-k test, so this kernel never executes in it.
The rest are fail-fast cancellations of sibling shards plus the pr-gate / *-finish aggregators.

Happy to dig into any of those if you would rather they were green first.

xiaofei-zheng · 2026-09-13T05:39:56Z

@HaiShaw thanks for the earlier approval. Could you help nudge this one along, or advise on how to proceed?

@DarkSharpness requested changes on Sep 9 with three performance comments. All three are addressed in 9a57b69a20 and measured on MI355X — the hot path is unchanged and the refinement path drops 29–35%. I replied with the numbers on Sep 10 and again on Sep 11 after a rebase, but there has been no response in four days, and the CHANGES_REQUESTED is still recorded against the old commit 7a83d7fdd5.

Every CI job that exercises this kernel is now green on both vendors, and mergeable is true again. Happy to keep waiting if a re-review is coming — just flagging that the requested work is done.

DarkSharpness · 2026-09-13T14:01:25Z

Hi. Could you please temporarily hold on until #38798 is merged? Sorry we're actively working on that and it should land within next few days. I will check this PR in detail later. This topk kernel is very sensitive in register usage on CUDA, and any seemingly irrelevant test might lead to wild performance data, so I will take a very detailed look at this PR later. Please stay tuned. Thanks :)

…verflowing threshold bin The coarse histogram bin is derived from the top bits of the fp16 cast of each score, so one bin spans about a quarter-binade. Phase 3 stages threshold-bin candidates into the fixed kMaxNumTie = 2048 staging buffer and phase 4 clips tie_count = min(count_eq, kMaxNumTie), so once a threshold bin holds more than 2048 elements handle_tie ranks whichever arrived first and the kernel returns a wrong selected-value multiset with no diagnostic. This is the shared DSA indexer top-k used by DeepSeek-V4.1 decode, prefill and the DeepGEMM candidate path, and it fails silently: every emitted index is a valid KV position, so nothing downstream detects the wrong selection. Refine the overflowing bin on the exact key instead of truncating it: up to three radix rounds (12/12/8 bits) over the order-preserving key from extract_exact_bin. It exits once the refined set fits the staging buffer, or once all 32 bits are consumed -- the key is injective on fp32 bit patterns, so the survivors are bit-identical by then and truncating them is legitimate. The fast path branches on count_eq > kMaxNumTie and is otherwise unchanged; kMaxNumTie, kBlockSize and the shared-memory footprint are untouched (the refinement histogram overlays the dead tie staging buffer in a union). This is a correctness fix, not a speedup: the hot path is unchanged. The CUDA-only TopKCluster path still truncates; refining there needs cluster-wide rather than block-local histogram and emit counters, and is left for a follow-up. Ported from sgl-project#37941, which does not apply cleanly after sgl-project#38829/sgl-project#39098 refactored this file. Identifier adaptations: smem->tie.handle -> smem->tie_handle, smem->tie.values -> smem->tie_values, smem->tie.refine_hist -> smem->refine_hist; the local sites name the counters count_gt/count_eq rather than above_count/equal_count. The PR's ROCm __launch_bounds__ hunk is intentionally dropped (CUDA-only deployment). Verified on H20 (SM90, cc 9.0): - The new test fails 8/10 on the unpatched kernel (the two all_equal controls pass, pinning the fault to truncation) and passes 10/10 patched. - test/registered/kernels/ops/attention/test_topk_v2.py: 298 passed. - cuobjdump resource usage and sizeof(Smem) are byte-identical to baseline: 31-32 registers, zero spills, shared memory unchanged.

xiaofei-zheng · 2026-09-14T06:06:33Z

Sounds good, we'll hold.

Two things worth flagging while we wait.

The truncation this PR fixes is unchanged in #38798 — the if (pos < kMaxNumTie) drop in the collect pass and the min(count_eq, kMaxNumTie) clip are both still there, on the register, streaming and cluster paths. So the kernel will still return a wrong top-k set when a threshold coarse bin overflows.

And please don't try to carry this change into #38798; it is safer for us to re-port it afterwards. The refinement re-reads candidates from global memory and relies on for_each_input visiting exactly the elements the collect pass counted. #38798 pads the last vector with padding_value() and compensates with atomicSub on the histogram, which breaks that invariant — easy to miss in a mechanical merge, and the failure mode would be a silently wrong refinement.

On your register point: we only have MI355X here, so every number in this PR is gfx950 (57 to 78 VGPRs, no spills, LDS unchanged). We cannot measure the CUDA side at all. If that is the main concern, it would help if someone with an NVIDIA box could check it — or tell us which shapes and metric you want and we will get them into the PR another way.

Ping me when #38798 lands and I'll rebase, re-measure and re-request review.

The coarse bin comes from the top bits of the fp16 cast, so one bin spans a quarter-binade. The collect pass stages threshold-bin candidates under `if (pos < kMaxNumTie)` and phase 4 clips `min(count_eq, kMaxNumTie)`, so an overflowing bin leaves `handle_tie` ranking an arrival-order subset and the kernel returns a wrong selected-value multiset with no diagnostic. Found while profiling GLM-5.2 decode on MI355X: device-side counters over an agentic serving run enter that path on ~0.04% of rows across all 8 ranks, with a largest observed bin of 4139 against the 2048 cap and none of those rows bit-identical. sgl-project#35257 reports a related problem from the CUDA side. sgl-project#39648 restructured this kernel but left all four truncation sites unchanged. Refine the overflowing bin on the exact key instead: up to three radix rounds over the order-preserving key from `extract_exact_bin`, 12 / 12 / 8 bits wide, exiting once the refined set fits the staging buffer or once all 32 bits are consumed, where the survivors are bit-identical and truncating them is exact. Each round emits its own "above" set while building the next round's histogram, so the refinement costs at most four input passes. The histogram overlays `tie_values`, which this path re-derives rather than trusting, so LDS is unchanged at 26928 B. The fast path branches on `count_eq > kMaxNumTie` and is otherwise untouched. Also pin the ROCm `__launch_bounds__` occupancy argument to the physical wave floor: HIP reads it as waves per SIMD where CUDA reads it as blocks per SM, so `kOccupancy` asked ROCm for 2 waves/SIMD, below the 4 a 1024-thread block already forces. The CUDA spelling stays because `topk_small_batch_cluster_kernel` shadows `kOccupancy` with a template parameter. Testing. Pre-existing cases draw from `torch.randn` and tolerate MAX_PERMIT_ERROR = 5, so they never reach the overflow. Added narrow, narrow_bin, tiny and two_values distributions that collapse a row into one coarse bin, plus overflow with DUAL_OUTPUT and overflow on an unaligned ragged window. On the kernel on main those 14 cases fail while the 2 all_equal controls pass, which pins the fault to truncation rather than to the binning. All 304 pass here. MI355X: the hot path is unchanged (5.324 vs 5.484 us at 6x8192, 12.094 vs 12.267 at 4x32768); rows that overflow now pay for work previously skipped by returning a wrong answer (narrow_bin 17.051 -> 21.684, two_values 36.194 -> 118.811 at 4x32768). 57 -> 79 VGPRs, zero spills. Known gap: the CUDA-only `TopKCluster` path still truncates. Refining there needs cluster-wide rather than block-local histogram and emit counters. Co-authored-by: Cursor <cursoragent@cursor.com>

xiaofei-zheng · 2026-09-16T09:45:05Z

Rebased onto #39648 and re-ported, now e503afa6bd.

A correction to what I said earlier: I claimed the padding change would break the refinement's counting invariant. That was wrong. padding_value() is quiet_NaN(), and NaN fails both val >= v_hi and val >= v_lo, so padding is counted by neither count_gt nor count_eq — the ragged path even writes the NaN into the score buffer itself. Since for_each_input visits exactly [0, seq_len), the collect pass and the refinement's re-read agree on which elements exist, and no extra handling was needed. There is now a test that pins this rather than leaving it to code reading: overflow on an unaligned ragged window, which fails on the kernel on main.

Also worth flagging: #39648 restructured the kernel but left all four truncation sites unchanged, so the defect is exactly as it was.

Verified on MI355X — 304 pass here, and against the kernel on main the 14 new overflow cases fail while the 2 all_equal controls pass. Numbers are in the updated description.

xiaofei-zheng requested review from BBuf, DarkSharpness, HaiShaw, HydraQYH, celve and yuan-luo as code owners September 4, 2026 06:24

github-actions Bot added the jit-kernel label Sep 4, 2026

xiaofei-zheng force-pushed the feature/xiaofei/topk-exact-threshold-bin branch 2 times, most recently from 8a746f3 to 3dbc304 Compare September 4, 2026 07:02

xiaofei-zheng changed the title ~~[deepseek_v4] Make the fused top-k exact on an overflowing threshold bin~~ [DSA] Make the fused top-k exact on an overflowing threshold bin Sep 4, 2026

xiaofei-zheng mentioned this pull request Sep 4, 2026

[DSA] Pin the ROCm top-k launch bound to the physical wave floor #37942

Closed

3 tasks

xiaofei-zheng force-pushed the feature/xiaofei/topk-exact-threshold-bin branch 3 times, most recently from ae5db33 to 7a83d7f Compare September 8, 2026 06:48

HaiShaw assigned 1am9trash Sep 9, 2026

HaiShaw approved these changes Sep 9, 2026

View reviewed changes

github-actions Bot added the run-ci label Sep 9, 2026

DarkSharpness requested changes Sep 9, 2026

View reviewed changes

xiaofei-zheng force-pushed the feature/xiaofei/topk-exact-threshold-bin branch from 7e0bc53 to c05fe16 Compare September 10, 2026 07:28

xiaofei-zheng force-pushed the feature/xiaofei/topk-exact-threshold-bin branch from c05fe16 to fcd3905 Compare September 10, 2026 10:53

xiaofei-zheng force-pushed the feature/xiaofei/topk-exact-threshold-bin branch from fcd3905 to 9a57b69 Compare September 12, 2026 01:39

xiaofei-zheng force-pushed the feature/xiaofei/topk-exact-threshold-bin branch from 9a57b69 to e503afa Compare September 16, 2026 09:44

Uh oh!

Conversation

xiaofei-zheng commented Sep 4, 2026 • edited by github-actions Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Modifications

Accuracy Tests

Benchmarking and Profiling

Known gap

Checklist

Known gap

CI States

Uh oh!

xiaofei-zheng commented Sep 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

HaiShaw commented Sep 9, 2026

Uh oh!

HaiShaw commented Sep 9, 2026

Uh oh!

DarkSharpness left a comment

Choose a reason for hiding this comment

Uh oh!

DarkSharpness Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

DarkSharpness Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

DarkSharpness Sep 9, 2026

Choose a reason for hiding this comment

Uh oh!

DarkSharpness left a comment

Choose a reason for hiding this comment

Uh oh!

xiaofei-zheng commented Sep 10, 2026

Uh oh!

xiaofei-zheng commented Sep 11, 2026

Uh oh!

xiaofei-zheng commented Sep 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

DarkSharpness commented Sep 13, 2026

Uh oh!

xiaofei-zheng commented Sep 14, 2026

Uh oh!

xiaofei-zheng commented Sep 16, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

xiaofei-zheng commented Sep 4, 2026 •

edited by github-actions Bot

Loading

xiaofei-zheng commented Sep 9, 2026 •

edited

Loading

xiaofei-zheng commented Sep 13, 2026 •

edited

Loading